/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#pragma once

#include <cuos/cuos.h>

//----------------------------------------------------------------------

static __device__ int open(cuos::device::sys_queue_t *sys_q, const char *glb_pathname, int flags)
{
    const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
    const uint block_size = blockDim.x*blockDim.y*blockDim.z;
    //const uint grid_size = gridDim.y*gridDim.x;
    const uint gid = tid + bid*block_size;
    int retcode = cuMPI_SUCCESS;

    using namespace cuos::system_service;

    if(bid == 0 && tid == 0) {
        cuos::sys_req_t sys_desc;
        
        sys_desc.id = sys_open;
        sys_desc.n_params = 2;
        sys_desc.flags = cuos::system_service::request::flag_pending;
        sys_desc.params[0] = cuos_make_ptr_param( glb_pathname, 1 );
        sys_desc.params[1] = cuos_make_int_param( flags );
        sys_desc.retval = cuos_make_int_param( 0 );
        
        //cuos_error("sys_q: initialized=%d req=%p rep=%p\n", 
        //           cuos::device::sys_q.m_initialized, cuos::device::sys_q.m_req, cuos::device::sys_q.m_rep);

        gbl_req_event.u.t = sys_desc;
    }
    __syncthreads();

    int ret = cuos::device::syscall_exec(sys_q, &gbl_req_event, &gbl_rep_event);
    if(ret != cuos::SUCCESS) {
        if(0 == gid)
            cuos_error("error in syscall_exec()\n");
        retcode = RUNTIME;
    } else
        retcode = gbl_rep_event.u.t.retval.si;

    return retcode;
}

//----------------------------------------------------------------------

static __device__ int close(cuos::device::sys_queue_t *sys_q, int glb_fd)
{
    const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
    const uint block_size = blockDim.x*blockDim.y*blockDim.z;
    //const uint grid_size = gridDim.y*gridDim.x;
    const uint gid = tid + bid*block_size;
    int retcode = cuMPI_SUCCESS;

    using namespace cuos::system_service;

    if(bid == 0 && tid == 0) {
        cuos::sys_req_t sys_desc;
        
        sys_desc.id = sys_close;
        sys_desc.n_params = 1;
        sys_desc.flags = cuos::system_service::request::flag_pending;
        sys_desc.params[0] = cuos_make_int_param( glb_fd );
        sys_desc.retval = cuos_make_int_param( 0 );
        
        //cuos_error("sys_q: initialized=%d req=%p rep=%p\n", 
        //           cuos::device::sys_q.m_initialized, cuos::device::sys_q.m_req, cuos::device::sys_q.m_rep);

        gbl_req_event.u.t = sys_desc;
    }
    __syncthreads();

    int ret = cuos::device::syscall_exec(sys_q, &gbl_req_event, &gbl_rep_event);
    if(ret != cuos::SUCCESS) {
        if(0 == gid)
            cuos_error("error in syscall_exec()\n");
        retcode = RUNTIME;
    } else
        retcode = gbl_rep_event.u.t.retval.si;

    return retcode;
}

//----------------------------------------------------------------------

static __device__ int read(cuos::device::sys_queue_t *sys_q, int glb_fd, void *glb_buf, size_t count)
{
    const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
    const uint block_size = blockDim.x*blockDim.y*blockDim.z;
    //const uint grid_size = gridDim.y*gridDim.x;
    const uint gid = tid + bid*block_size;
    int retcode = cuMPI_SUCCESS;

    using namespace cuos::system_service;

    if(bid == 0 && tid == 0) {
        cuos::sys_req_t sys_desc;
        
        sys_desc.id = sys_read;
        sys_desc.n_params = 3;
        sys_desc.flags = cuos::system_service::request::flag_pending;
        sys_desc.params[0] = cuos_make_int_param( glb_fd, );
        sys_desc.params[1] = cuos_make_ptr_param( glb_buf, 1 );
        sys_desc.params[2] = cuos_make_uint_param( count );
        sys_desc.retval = cuos_make_int_param( 0 );
        
        //cuos_error("sys_q: initialized=%d req=%p rep=%p\n", 
        //           cuos::device::sys_q.m_initialized, cuos::device::sys_q.m_req, cuos::device::sys_q.m_rep);

        gbl_req_event.u.t = sys_desc;
    }
    __syncthreads();

    int ret = cuos::device::syscall_exec(sys_q, &gbl_req_event, &gbl_rep_event);
    if(ret != cuos::SUCCESS) {
        if(0 == gid)
            cuos_error("error in syscall_exec()\n");
        retcode = RUNTIME;
    } else
        retcode = gbl_rep_event.u.t.retval.si;

    return retcode;
}

//----------------------------------------------------------------------

static __device__ int write(cuos::device::sys_queue_t *sys_q, int glb_fd, const void *glb_buf, size_t count)
{
    const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
    const uint block_size = blockDim.x*blockDim.y*blockDim.z;
    //const uint grid_size = gridDim.y*gridDim.x;
    const uint gid = tid + bid*block_size;
    int retcode = cuMPI_SUCCESS;

    using namespace cuos::system_service;

    if(bid == 0 && tid == 0) {
        cuos::sys_req_t sys_desc;
        
        sys_desc.id = sys_write;
        sys_desc.n_params = 3;
        sys_desc.flags = cuos::system_service::request::flag_pending;
        sys_desc.params[0] = cuos_make_int_param( glb_fd, );
        sys_desc.params[1] = cuos_make_ptr_param( glb_buf, 1 );
        sys_desc.params[2] = cuos_make_uint_param( count );
        sys_desc.retval = cuos_make_int_param( 0 );
        
        //cuos_error("sys_q: initialized=%d req=%p rep=%p\n", 
        //           cuos::device::sys_q.m_initialized, cuos::device::sys_q.m_req, cuos::device::sys_q.m_rep);

        gbl_req_event.u.t = sys_desc;
    }
    __syncthreads();

    int ret = cuos::device::syscall_exec(sys_q, &gbl_req_event, &gbl_rep_event);
    if(ret != cuos::SUCCESS) {
        if(0 == gid)
            cuos_error("error in syscall_exec()\n");
        retcode = RUNTIME;
    } else
        retcode = gbl_rep_event.u.t.retval.si;

    return retcode;
}

//----------------------------------------------------------------------

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
